


#define reconstructType unsigned char*



//median required for the MEDIAN vector predictor
__device__ int cuda_median_3(int a, int b, int c)
{
    if(a>b){
        if(c>b){
            if(c>a) b=a;
            else    b=c;
        }
    }else{
        if(b>c){
            if(c>a) b=c;
            else    b=a;
        }
    }
    return b;
}


#define MBLOCKS_PER_MSWARP 4


__device__ void half_pixel_warp_sad( int half_x, int half_y,unsigned char* my_part_of_current_block,const unsigned char* __restrict__ reference_mblock, int* shared_distortions, const int internal_tx) {

    int pix_x = (8*internal_tx) % 16;
    int pix_y = (8*internal_tx) / 16;
    
    if (half_x && !half_y) {
        if (half_x + pix_x > 0) {
        shared_distortions[internal_tx] = abs(my_part_of_current_block[0] 
                                            - ((reference_mblock[ pix_x + pix_y*16] + reference_mblock[ pix_x + pix_y*16 + half_x] +1)>>1)  );
        }
        else {
        shared_distortions[internal_tx] = 0;
        }
        shared_distortions[internal_tx] += abs(my_part_of_current_block[1]
                            - ((reference_mblock[ pix_x + 1 + pix_y*16] + reference_mblock[ pix_x + 1 +pix_y*16 + half_x] +1)>>1)  );
                            
        shared_distortions[internal_tx] += abs(my_part_of_current_block[2] 
                            - ((reference_mblock[ pix_x + 2 + pix_y*16] + reference_mblock[ pix_x + 2 +pix_y*16 + half_x] +1)>>1)  );
                   
        shared_distortions[internal_tx] += abs(my_part_of_current_block[3]  
        - ((reference_mblock[ pix_x + 3 + pix_y*16] + reference_mblock[ pix_x + 3 +pix_y*16 + half_x] +1)>>1)  );
        
        shared_distortions[internal_tx] += abs(my_part_of_current_block[4]  
        - ((reference_mblock[ pix_x + 4 + pix_y*16] + reference_mblock[ pix_x + 4 +pix_y*16 + half_x] +1)>>1)  );
        
        shared_distortions[internal_tx] += abs(my_part_of_current_block[5]  
        - ((reference_mblock[ pix_x + 5 + pix_y*16] + reference_mblock[ pix_x + 5 +pix_y*16 + half_x] +1)>>1)  );
        
        shared_distortions[internal_tx] += abs(my_part_of_current_block[6]  
        - ((reference_mblock[ pix_x + 6 + pix_y*16] + reference_mblock[ pix_x + 6 +pix_y*16 + half_x] +1)>>1)  );
        
        if ((half_x + pix_x  + 7) < 16) {
            shared_distortions[internal_tx] += abs(my_part_of_current_block[7]  
            - ((reference_mblock[ pix_x + 7 + pix_y*16] + reference_mblock[ pix_x + 7 +pix_y*16 + half_x] +1)>>1)  );
        }
        
   }
   else if ( half_y && (!half_x) ) {
    if ((half_y + pix_y > 0) && (half_y + pix_y < 16)) {
        shared_distortions[internal_tx] = abs(my_part_of_current_block[0] 
                                                        - (( reference_mblock[ pix_x + pix_y*16] + reference_mblock[ pix_x + (pix_y+half_y)*16] + 1)>>1) );
                                                        
        shared_distortions[internal_tx] += abs(my_part_of_current_block[1] 
                                                    - ( (reference_mblock[ pix_x + 1 +pix_y*16] + reference_mblock[ pix_x + 1 +(pix_y+half_y)*16] + 1)>>1) );
                                                    
        shared_distortions[internal_tx] += abs(my_part_of_current_block[2] 
                                                    - (( reference_mblock[ pix_x + 2 + pix_y*16] + reference_mblock[ pix_x + 2+ (pix_y+half_y)*16] + 1)>>1) );
                                                    
        shared_distortions[internal_tx] += abs(my_part_of_current_block[3] 
                                            - (( reference_mblock[ pix_x + 3 + pix_y*16] + reference_mblock[ pix_x + 3+ (pix_y+half_y)*16] + 1)>>1) );
                                            
        shared_distortions[internal_tx] += abs(my_part_of_current_block[4] 
                                            - (( reference_mblock[ pix_x + 4 + pix_y*16] + reference_mblock[ pix_x + 4+ (pix_y+half_y)*16] + 1)>>1) );
                                            
        shared_distortions[internal_tx] += abs(my_part_of_current_block[5] 
                                            - (( reference_mblock[ pix_x + 5 + pix_y*16] + reference_mblock[ pix_x + 5+ (pix_y+half_y)*16] + 1)>>1) );
                                            
        shared_distortions[internal_tx] += abs(my_part_of_current_block[6] 
                                            - (( reference_mblock[ pix_x + 6 + pix_y*16] + reference_mblock[ pix_x + 6+ (pix_y+half_y)*16] + 1)>>1) );                                            
        shared_distortions[internal_tx] += abs(my_part_of_current_block[7] 
                                            - (( reference_mblock[ pix_x + 7 + pix_y*16] + reference_mblock[ pix_x + 7+ (pix_y+half_y)*16] + 1)>>1) );                                            

       }
       else {
        shared_distortions[internal_tx] = 0;
       }
   } else {
        ///hacer la busqueda half pixel en diagonal
        if ((half_y + pix_y > 0) && (half_y + pix_y < 16)) {
            if (half_x + pix_x > 0) {
                shared_distortions[internal_tx] = abs(my_part_of_current_block[0] 
                                            - ((reference_mblock[ pix_x + pix_y*16] + reference_mblock[ pix_x + (pix_y+ half_y)*16 + half_x] + 
                                            + reference_mblock[ pix_x + (pix_y+ half_y)*16 ] + reference_mblock[ pix_x + (pix_y)*16 + half_x ] + 2)>>2)  );
            }
            else {
                shared_distortions[internal_tx] = 0;
            }
            shared_distortions[internal_tx] += abs(my_part_of_current_block[1]
                            - ((reference_mblock[ pix_x + 1 + pix_y*16] + reference_mblock[ pix_x + 1 + (pix_y+ half_y)*16 + half_x]
                             + reference_mblock[ pix_x + 1 + (pix_y+ half_y)*16 ] + + reference_mblock[ pix_x + 1 + (pix_y)*16 + half_x] +2)>>2)  );
                             
            shared_distortions[internal_tx] += abs(my_part_of_current_block[2] 
                            - ((reference_mblock[ pix_x + 2 + pix_y*16] + reference_mblock[ pix_x + 2 + (pix_y+ half_y)*16 + half_x] 
                            + reference_mblock[ pix_x + 2 + (pix_y+ half_y)*16 ] + reference_mblock[ pix_x + 2 + (pix_y)*16 + half_x] +2)>>2)  );

            shared_distortions[internal_tx] += abs(my_part_of_current_block[3] 
                            - ((reference_mblock[ pix_x + 3 + pix_y*16] + reference_mblock[ pix_x + 3 + (pix_y+ half_y)*16 + half_x] 
                            + reference_mblock[ pix_x + 3 + (pix_y+ half_y)*16 ] + reference_mblock[ pix_x + 3 + (pix_y)*16 + half_x] +2)>>2)  );

            shared_distortions[internal_tx] += abs(my_part_of_current_block[4] 
                            - ((reference_mblock[ pix_x + 4 + pix_y*16] + reference_mblock[ pix_x + 4 + (pix_y+ half_y)*16 + half_x] 
                            + reference_mblock[ pix_x + 4 + (pix_y+ half_y)*16 ] + reference_mblock[ pix_x + 4 + (pix_y)*16 + half_x] +2)>>2)  );

            shared_distortions[internal_tx] += abs(my_part_of_current_block[5] 
                            - ((reference_mblock[ pix_x + 5 + pix_y*16] + reference_mblock[ pix_x + 5 + (pix_y+ half_y)*16 + half_x] 
                            + reference_mblock[ pix_x + 5 + (pix_y+ half_y)*16 ] + reference_mblock[ pix_x + 5 + (pix_y)*16 + half_x] +2)>>2)  );

            shared_distortions[internal_tx] += abs(my_part_of_current_block[6] 
                            - ((reference_mblock[ pix_x + 6 + pix_y*16] + reference_mblock[ pix_x + 6 + (pix_y+ half_y)*16 + half_x] 
                            + reference_mblock[ pix_x + 6 + (pix_y+ half_y)*16 ] + reference_mblock[ pix_x + 6 + (pix_y)*16 + half_x] +2)>>2)  );
                                                        
            if ((half_x + pix_x+ 7) < 16) {                   
                shared_distortions[internal_tx] += abs(my_part_of_current_block[7]  
                - ((reference_mblock[ pix_x + 7 + pix_y*16] + reference_mblock[ pix_x + 7 + (pix_y+ half_y)*16 + half_x] 
                + reference_mblock[ pix_x + 7 + (pix_y+ half_y)*16 ] + reference_mblock[ pix_x + 7 + (pix_y)*16 + half_x] +2)>>2)  );
            }
        }
        else {
            shared_distortions[internal_tx] = 0;
        }

   }

    if ( (internal_tx % 2)==0) {
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 1];
    }

    if ( (internal_tx % 4)==0) {
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 2];
    }

    if ( (internal_tx % 8)==0) {
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 4];
    }

    if ((internal_tx % 16)==0) {
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 8];
    }


    if ( (internal_tx % 32)==0) {
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 16];
        shared_distortions[internal_tx] = ((16 + (half_x && half_y))*shared_distortions[0])/(15);
    }



}

//the current frame to encode is stored in the texture srcTex
texture<uchar4, 2, cudaReadModeElementType> srcTex;
//we store the vector costs for RDO
texture<int, 2, cudaReadModeElementType> RDO_table;

#define RDO_DISPLACEMENT_SHIFT 128

__global__ void motion_search_granularity(reconstructType reference, unsigned int referenceStride, unsigned int* vectors, unsigned int vectorsStride, int frameWidth, int frameHeight, reconstructType rawluma_out) {

    const int mb_x = MBLOCKS_PER_MSWARP*blockIdx.x + threadIdx.x/32;
    const int mb_y = blockIdx.y;
    
    //this is to simulate a 32-thread group within the cuda block (which is larger for performance issues). This thread group will target one macroblock
    const int internal_tx = threadIdx.x % 32;
          
    int pix_x = (8*internal_tx) % 16; //each thread reads 8 pixels in the same row
    int pix_y = (8*internal_tx) / 16; //each 2 threads read one row
    
    unsigned char my_part_of_current_block[8];
    
    //LOADING THE CURRENT RAW FRAME AND CALCULATING THE LUMA COMPONENT
    uchar4 color = tex2D(srcTex, 16*mb_x + pix_x, frameHeight - 16*mb_y - pix_y -1);
//	64, 129, 25, 128, 8, 16
    uint8_t ctrans = ((64*color.x + 129*color.y + 25*color.z + 128)>>8) + 16; //luma
    my_part_of_current_block[0] = ctrans; // <- each threads loads 8 pixels row-wise
    
    color = tex2D(srcTex, 16*mb_x + pix_x +1, frameHeight - 16*mb_y - pix_y -1);
    ctrans = ((64*color.x + 129*color.y + 25*color.z + 128)>>8) + 16; //luma
    my_part_of_current_block[1] = ctrans;
    
    color = tex2D(srcTex, 16*mb_x + pix_x +2, frameHeight - 16*mb_y - pix_y -1);
    ctrans = ((64*color.x + 129*color.y + 25*color.z + 128)>>8) + 16; //luma
    my_part_of_current_block[2] = ctrans;
    
    color = tex2D(srcTex, 16*mb_x + pix_x +3, frameHeight - 16*mb_y - pix_y -1);
    ctrans = ((64*color.x + 129*color.y + 25*color.z + 128)>>8) + 16; //luma
    my_part_of_current_block[3] = ctrans;
    
    color = tex2D(srcTex, 16*mb_x + pix_x + 4, frameHeight - 16*mb_y - pix_y -1);
    ctrans = ((64*color.x + 129*color.y + 25*color.z + 128)>>8) + 16; //luma
    my_part_of_current_block[4] = ctrans;
    
    color = tex2D(srcTex, 16*mb_x + pix_x + 5, frameHeight - 16*mb_y - pix_y -1);
    ctrans = ((64*color.x + 129*color.y + 25*color.z + 128)>>8) + 16; //luma
    my_part_of_current_block[5] = ctrans;
    
    color = tex2D(srcTex, 16*mb_x + pix_x + 6, frameHeight - 16*mb_y - pix_y -1);
    ctrans = ((64*color.x + 129*color.y + 25*color.z + 128)>>8) + 16; //luma
    my_part_of_current_block[6] = ctrans;
    
    color = tex2D(srcTex, 16*mb_x + pix_x + 7, frameHeight - 16*mb_y - pix_y -1);
    ctrans = ((64*color.x + 129*color.y + 25*color.z + 128)>>8) + 16; //luma
    my_part_of_current_block[7] = ctrans;  
    
     //we write the 8 pixels of the LUMA transform for further use
    *((long*) (rawluma_out + 16*mb_x + pix_x + (16*mb_y + pix_y)*referenceStride)) = *((long*)my_part_of_current_block);
    
 
           
    __shared__ unsigned char __reference_mblock[256*MBLOCKS_PER_MSWARP];
    unsigned char* reference_mblock = __reference_mblock + 256*(threadIdx.x/32);
    __shared__ int __shared_distortions[ MBLOCKS_PER_MSWARP*32 ];
    int* shared_distortions = __shared_distortions + 32*(threadIdx.x/32);
    int best_vx = 0;
    int best_vy = 0;

    int best_distortion = 99999;
    
    int current_vx = 0;
    int current_vy = 0;
    int continue_sad = 0;
    
    //macro for testing if we improve the best soultion with the 
    #define TEST_MOT_VECT_MACRO_3() \
    SAD_MACRO_3()\
    if (shared_distortions[0] < best_distortion ) {\
            best_distortion = shared_distortions[0];\
            best_vx = current_vx;\
            best_vy = current_vy;\
            continue_sad = 1; \
     }\

    
    const int coord_x = 16*mb_x;
    const int coord_y = 16*mb_y;
    
    int we_sad = 0;
#define SAD_MACRO_3() \
    we_sad = ((coord_x + current_vx) > -1) && ((coord_x + current_vx + 15) < frameWidth) && ((coord_y + current_vy) > -1) && ((coord_y + current_vy + 15) < frameHeight);\
     \
    if (we_sad) { \
    pix_x = internal_tx % 16;\
    pix_y = internal_tx / 16;\
\
    \
    reference_mblock[ pix_x + pix_y*16] = *(reference +coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y)*referenceStride);\
    reference_mblock[ pix_x + (pix_y + 2)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 2)*referenceStride);\
    reference_mblock[ pix_x + (pix_y + 4)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 4)*referenceStride);\
    reference_mblock[ pix_x + (pix_y + 6)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 6)*referenceStride);\
    reference_mblock[ pix_x + (pix_y + 8)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 8)*referenceStride);\
    reference_mblock[ pix_x + (pix_y + 10)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 10)*referenceStride);\
    reference_mblock[ pix_x + (pix_y + 12)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 12)*referenceStride);\
    reference_mblock[ pix_x + (pix_y + 14)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 14)*referenceStride);\
    pix_x = (8*internal_tx) % 16;\
    pix_y = (8*internal_tx) / 16;\
    \
    shared_distortions[internal_tx] = abs(my_part_of_current_block[0] - reference_mblock[ pix_x + pix_y*16] );\
    shared_distortions[internal_tx] += abs(my_part_of_current_block[1] - reference_mblock[ pix_x +1 + pix_y*16] );\
    shared_distortions[internal_tx] += abs(my_part_of_current_block[2] - reference_mblock[ pix_x +2 + pix_y*16] );\
    shared_distortions[internal_tx] += abs(my_part_of_current_block[3] - reference_mblock[ pix_x  +3 + pix_y*16] );\
    shared_distortions[internal_tx] += abs(my_part_of_current_block[4] - reference_mblock[ pix_x  +4 + pix_y*16] );\
    shared_distortions[internal_tx] += abs(my_part_of_current_block[5] - reference_mblock[ pix_x  +5 + pix_y*16] );\
    shared_distortions[internal_tx] += abs(my_part_of_current_block[6] - reference_mblock[ pix_x  +6 + pix_y*16] );\
    shared_distortions[internal_tx] += abs(my_part_of_current_block[7] - reference_mblock[ pix_x  +7 + pix_y*16] );\
    \
    shared_distortions[internal_tx] += (!we_sad)*25; \
\
    \
    if ((internal_tx % 2)==0) {\
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 1];\
    }\
    if ((internal_tx % 4)==0) {\
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 2];\
    }\
    if ((internal_tx % 8)==0) {\
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 4];\
    }\
    if ((internal_tx % 16)==0) {\
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 8];\
    }\
    if ((internal_tx % 32)==0) {\
        shared_distortions[internal_tx] += shared_distortions[internal_tx + 16] + tex2D(RDO_table, RDO_DISPLACEMENT_SHIFT + current_vx - mvp_x, RDO_DISPLACEMENT_SHIFT + current_vy - mvp_y);\
    }\
    \
    }\
    else {\
        if (internal_tx ==0) { \
            shared_distortions[0] = 99999;\
        }    \
    }\


   
   #define PREDICTORS_ARE_DIFFERENT( pred1, pred2) (pred1.x != pred2.x || pred1.y != pred2.y)    
   
   __shared__ int _predictor_vectors[ 12 * MBLOCKS_PER_MSWARP];
   int* predictor_vectors = _predictor_vectors + 12 * (threadIdx.x/32);
   
   #define CHECK_DIFF_VECT( num_vect ) ((predictor_vectors[2*num_vect] != current_vx) | (predictor_vectors[2*num_vect +1 ] != current_vy))
   #define SAVE_PRED_VECT( num_vect ) { \
        predictor_vectors[2*num_vect] = current_vx; \
        predictor_vectors[2*num_vect +1 ] = current_vy; \
    }
   
   shared_distortions[internal_tx] = 66666;


    int mvp_x, mvp_y;
   //load and check the first predictor (collocated on the previous frame), only check if it has not been checked before
   unsigned int vect_code = *(vectors + mb_x + (mb_y)*vectorsStride) ;
   current_vx = ((vect_code & 0x0000FFFF) - 256)>>1;
   current_vy = ((vect_code >> 16) - 256)>>1;
   mvp_x = current_vx;
   mvp_y = current_vy;
   if ( current_vx | current_vy) {
    SAVE_PRED_VECT(0);
    TEST_MOT_VECT_MACRO_3();
   }
   


   	//test the 0,0 vector
    current_vx = 0;
    current_vy = 0;
    TEST_MOT_VECT_MACRO_3();
    if (best_distortion < 512) { //early stop
        if(internal_tx == 0) {
            *(vectors + mb_x + mb_y*vectorsStride) = ( (2*best_vy + 256) << 16 ) | (2*best_vx +  256);
        }
        return;
    }
    
   

   //top predictor on the previous frame (or on this it has already been written on the previous  frame
   if (mb_y) {
   unsigned int vect_code = *(vectors + mb_x + (mb_y-1)*vectorsStride) ;
   current_vx = ((vect_code & 0x0000FFFF) - 256)>>1;
   current_vy = ((vect_code >> 16) - 256)>>1;
   if ( (current_vx | current_vy) && CHECK_DIFF_VECT(0) ) {
    SAVE_PRED_VECT(1);
    TEST_MOT_VECT_MACRO_3();
   }
   }
   //left predictor on the previous frame
   if (mb_x) {
   unsigned int vect_code = *(vectors + (mb_x-1) + (mb_y)*vectorsStride) ;
   current_vx = ((vect_code & 0x0000FFFF) - 256)>>1;
   current_vy = ((vect_code >> 16) - 256)>>1;
   if ( (current_vx | current_vy) && CHECK_DIFF_VECT(0) && CHECK_DIFF_VECT(1) ) {
    SAVE_PRED_VECT(2);
    TEST_MOT_VECT_MACRO_3();
   }
   }
   //median of the top, left, and topright
   if ( ((mb_x*16 + 16) < (frameWidth)) && mb_x && mb_y) {
       unsigned int vect_code = *(vectors + (mb_x-1) + (mb_y)*vectorsStride) ;
       int current_vx_l = ((vect_code & 0x0000FFFF) - 256)>>1;
       int current_vy_l = ((vect_code >> 16) - 256)>>1;
       vect_code = *(vectors + (mb_x) + (mb_y-1)*vectorsStride) ;
       int current_vx_t = ((vect_code & 0x0000FFFF) - 256)>>1;
       int current_vy_t = ((vect_code >> 16) - 256)>>1;
       vect_code = *(vectors + (mb_x+1) + (mb_y-1)*vectorsStride) ;
       int current_vx_tr = ((vect_code & 0x0000FFFF) - 256)>>1;
       int current_vy_tr = ((vect_code >> 16) - 256)>>1;
       current_vx = cuda_median_3(current_vx_l , current_vx_t, current_vx_tr); 
       current_vy = cuda_median_3(current_vy_l , current_vy_t, current_vy_tr);

       if (abs(current_vx) > 128) current_vx = 0;
       if (abs(current_vy) > 128) current_vy = 0;
       if ( (current_vx | current_vy) && CHECK_DIFF_VECT(0) && CHECK_DIFF_VECT(1) && CHECK_DIFF_VECT(2) ){
       SAVE_PRED_VECT(3);
       TEST_MOT_VECT_MACRO_3();
       }
   }


    //right predictor
   if ( (coord_x + 16) < (frameWidth) ) {
       unsigned int vect_code = *(vectors + (mb_x+1) + (mb_y)*vectorsStride) ;
       current_vx = ((vect_code & 0x0000FFFF) - 256)>>1;
       current_vy = ((vect_code >> 16) - 256)>>1;
       if (abs(current_vx) > 256) current_vx = 0;
       if (abs(current_vy) > 256) current_vy = 0;
       if ( (current_vx | current_vy) && CHECK_DIFF_VECT(0) && CHECK_DIFF_VECT(1) && CHECK_DIFF_VECT(2) && CHECK_DIFF_VECT(3)){
       SAVE_PRED_VECT(4);
       TEST_MOT_VECT_MACRO_3();
       }
   }

  //down predictor
    if ((coord_y + 16) < frameHeight) {
       unsigned int vect_code = *(vectors + (mb_x) + (mb_y+1)*vectorsStride) ;
       current_vx = ((vect_code & 0x0000FFFF) - 256)>>1;
       current_vy = ((vect_code >> 16) - 256)>>1;
       if (abs(current_vx) > 256) current_vx = 0;
       if (abs(current_vy) > 256) current_vy = 0;
       if ( (current_vx | current_vy) && CHECK_DIFF_VECT(0) && CHECK_DIFF_VECT(1) && CHECK_DIFF_VECT(2) && CHECK_DIFF_VECT(3) && CHECK_DIFF_VECT(4)){
       TEST_MOT_VECT_MACRO_3();
       }
   }
   

    int center_vx = 0;
    int center_vy = 0;

    //small diamond search   
    continue_sad = 0;
    center_vx = best_vx;
    center_vy = best_vy;
    current_vx = center_vx + 1;
    current_vy = center_vy + 0;
    TEST_MOT_VECT_MACRO_3()
    current_vx = center_vx + -1;
    current_vy = center_vy + 0;
    TEST_MOT_VECT_MACRO_3()
    current_vx = center_vx + 0;
    current_vy = center_vy + 1; 
    TEST_MOT_VECT_MACRO_3()
    current_vx = center_vx + 0;
    current_vy = center_vy + -1; 
    TEST_MOT_VECT_MACRO_3()          
    
    
    while ( continue_sad) {
    continue_sad = 0;
    int delta_x = best_vx - center_vx;
    
    center_vx = best_vx;
    center_vy = best_vy;
    
    if (delta_x) {
        current_vx = center_vx + delta_x;
        current_vy = center_vy + 0;
        TEST_MOT_VECT_MACRO_3()
        current_vx = center_vx + 0;
        current_vy = center_vy + 1; 
        TEST_MOT_VECT_MACRO_3()
        current_vx = center_vx + 0;
        current_vy = center_vy + -1; 
        TEST_MOT_VECT_MACRO_3()
    }
    else {
        int delta_y = best_vy - center_vy;
        current_vx = center_vx + 1;
        current_vy = center_vy + 0;
        TEST_MOT_VECT_MACRO_3()
        current_vx = center_vx + -1;
        current_vy = center_vy + 0;
        TEST_MOT_VECT_MACRO_3()
        current_vx = center_vx + 0;
        current_vy = center_vy + delta_y;
        TEST_MOT_VECT_MACRO_3()
        }
   } 

   int best_half_x = 0 ;
   int best_half_y = 0;
   
   int half_x = 1;
   int half_y = 0;
   
   pix_x = internal_tx % 16;
   pix_y = internal_tx / 16;
   current_vx = best_vx;
   current_vy = best_vy;

   //reload the best mb to calc half pixel accuracy
   reference_mblock[ pix_x + pix_y*16] = *(reference +coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y)*referenceStride);
   reference_mblock[ pix_x + (pix_y + 2)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 2)*referenceStride);
   reference_mblock[ pix_x + (pix_y + 4)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 4)*referenceStride);
   reference_mblock[ pix_x + (pix_y + 6)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 6)*referenceStride);
   reference_mblock[ pix_x + (pix_y + 8)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 8)*referenceStride);
   reference_mblock[ pix_x + (pix_y + 10)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 10)*referenceStride);
   reference_mblock[ pix_x + (pix_y + 12)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 12)*referenceStride);
   reference_mblock[ pix_x + (pix_y + 14)*16] = *(reference + coord_x + current_vx + pix_x +  (coord_y + current_vy + pix_y + 14)*referenceStride);

   if ( ((coord_x + best_vx + half_x) > -1) && ((coord_x + best_vx + half_x + 15) < frameWidth) && ((coord_y + best_vy + half_y) > -1) && ((coord_y + best_vy + half_y +15) < frameHeight) ) {
   half_pixel_warp_sad( half_x, half_y, my_part_of_current_block, reference_mblock, shared_distortions, internal_tx);

    if (shared_distortions[0] < best_distortion ) {
            best_distortion = shared_distortions[0];
            best_half_x = half_x ;
            best_half_y = half_y;
    }
    }

    half_x = -1;
    half_y = 0;

   if ( ((coord_x + best_vx + half_x) > -1) && ((coord_x + best_vx + half_x + 15) < frameWidth) && ((coord_y + best_vy + half_y) > -1) && ((coord_y + best_vy + half_y +15) < frameHeight) ) {
    half_pixel_warp_sad( half_x, half_y, my_part_of_current_block, reference_mblock, shared_distortions, internal_tx);
    if (shared_distortions[0] < best_distortion ) {
            best_distortion = shared_distortions[0];
            best_half_x = half_x ;
            best_half_y = half_y;
    }
    }

    half_x = 0;
    half_y = 1;
    
   if ( ((coord_x + best_vx + half_x) > -1) && ((coord_x + best_vx + half_x + 15) < frameWidth) && ((coord_y + best_vy + half_y) > -1) && ((coord_y + best_vy + half_y +15) < frameHeight) ) {
    half_pixel_warp_sad( half_x, half_y, my_part_of_current_block, reference_mblock, shared_distortions, internal_tx);
    if (shared_distortions[0] < best_distortion ) {
            best_distortion = shared_distortions[0];
            best_half_x = half_x ;
            best_half_y = half_y;
    }
    }

    half_x = 0;
    half_y = -1;
    
   if ( ((coord_x + best_vx + half_x) > -1) && ((coord_x + best_vx + half_x + 15) < frameWidth) && ((coord_y + best_vy + half_y) > -1) && ((coord_y + best_vy + half_y +15) < frameHeight) ) {
    half_pixel_warp_sad( half_x, half_y, my_part_of_current_block, reference_mblock, shared_distortions, internal_tx);
    if (shared_distortions[0] < best_distortion ) {
            best_distortion = shared_distortions[0];
            best_half_x = half_x ;
            best_half_y = half_y;
    }
   }

    half_x = -1;
    half_y = -1;
    
   if ( ((coord_x + best_vx + half_x) > -1) && ((coord_x + best_vx + half_x + 15) < frameWidth) && ((coord_y + best_vy + half_y) > -1) && ((coord_y + best_vy + half_y +15) < frameHeight) ) {
    half_pixel_warp_sad( half_x, half_y, my_part_of_current_block, reference_mblock, shared_distortions, internal_tx);
    if (shared_distortions[0] < best_distortion ) {
            best_distortion = shared_distortions[0];
            best_half_x = half_x ;
            best_half_y = half_y;
    }
   }

    half_x = 1;
    half_y = -1;
    
   if ( ((coord_x + best_vx + half_x) > -1) && ((coord_x + best_vx + half_x + 15) < frameWidth) && ((coord_y + best_vy + half_y) > -1) && ((coord_y + best_vy + half_y +16) < frameHeight) ) {
    half_pixel_warp_sad( half_x, half_y, my_part_of_current_block, reference_mblock, shared_distortions, internal_tx);
    if (shared_distortions[0] < best_distortion ) {
            best_distortion = shared_distortions[0];
            best_half_x = half_x ;
            best_half_y = half_y;
    }
   }

    half_x = 1;
    half_y = 1;
    
   if ( ((coord_x + best_vx + half_x) > -1) && ((coord_x + best_vx + half_x + 15) < frameWidth) && ((coord_y + best_vy + half_y) > -1) && ((coord_y + best_vy + half_y +15) < frameHeight) ) {
    half_pixel_warp_sad( half_x, half_y, my_part_of_current_block, reference_mblock, shared_distortions, internal_tx);
    if (shared_distortions[0] < best_distortion ) {
            best_distortion = shared_distortions[0];
            best_half_x = half_x ;
            best_half_y = half_y;
    }
   }

    half_x = -1;
    half_y = 1;
    
   if ( ((coord_x + best_vx + half_x) > -1) && ((coord_x + best_vx + half_x + 15) < frameWidth) && ((coord_y + best_vy + half_y) > -1) && ((coord_y + best_vy + half_y +15) < frameHeight) ) {
    half_pixel_warp_sad( half_x, half_y, my_part_of_current_block, reference_mblock, shared_distortions, internal_tx);
    if (shared_distortions[0] < best_distortion ) {
            best_distortion = shared_distortions[0];
            best_half_x = half_x ;
            best_half_y = half_y;
    }
   }
   

    if (internal_tx==0) {
    //best_vx = 0;best_half_x = 0;
    //best_vy = 0;best_half_y = 0;
      //  best_vx = 0;
      //  best_vy = 0;

      if (abs(best_vx ) > 126) {best_vx = 0;  best_half_x = 0;}
      if (abs(best_vy ) > 126) {best_vy = 0; best_half_y = 0;}
      
        *(vectors + mb_x + mb_y*vectorsStride) = ( (2*best_vy + best_half_y+ 256) << 16 ) | (2*best_vx + best_half_x + 256);
    }
   // __syncthreads();

}

